//
//  ParticleSystem.m
//  DeferredLighting
//
//  Created by Holmes Futrell on 2/14/11.
//  Copyright 2011 __MyCompanyName__. All rights reserved.
//

#import "ParticleSystem.h"
#import "rb_error_checking.h"
#import "clb_error.h"
#import "clb_event_profiling.h"

@implementation ParticleSystem

@synthesize numParticles, eyespaceParticles;

-(id)initWithNumParticles:(int)_numParticles simpleCompute:(SimpleCompute *)_simpleCompute {
	
	self = [super init];
	
	numParticles = _numParticles;
	
	simpleCompute = [_simpleCompute retain];
	
	glGenBuffers(1, &glPositionsBuffer);
	glBindBuffer(GL_ARRAY_BUFFER, glPositionsBuffer);
	glBufferData(GL_ARRAY_BUFFER, sizeof(float) * 4 * numParticles, NULL, GL_DYNAMIC_DRAW);
	
	glGenBuffers(1, &glVelocitiesBuffer);
	glBindBuffer(GL_ARRAY_BUFFER, glVelocitiesBuffer);
	glBufferData(GL_ARRAY_BUFFER, sizeof(float) * 4 * numParticles, NULL, GL_DYNAMIC_DRAW);
	
	glGenBuffers(1, &glAgeBuffer);
	glBindBuffer(GL_ARRAY_BUFFER, glAgeBuffer);
	glBufferData(GL_ARRAY_BUFFER, sizeof(float) * 1 * numParticles, NULL, GL_DYNAMIC_DRAW);
	
	
	rb_check_error_simple();
	
	cl_int error;
	
	clPositionsBuffer = clCreateFromGLBuffer(simpleCompute.context, CL_MEM_READ_WRITE, glPositionsBuffer, &error);
	clb_check_error(error, "clCreateFromGLBuffer");
	clVelocitiesBuffer = clCreateFromGLBuffer(simpleCompute.context, CL_MEM_READ_WRITE, glVelocitiesBuffer, &error);
	clb_check_error(error, "clCreateFromGLBuffer");
	clAgeBuffer = clCreateFromGLBuffer(simpleCompute.context, CL_MEM_READ_WRITE, glAgeBuffer, &error);
	clb_check_error(error, "clCreateFromGLBuffer");


	eyespaceParticles = clCreateBuffer(simpleCompute.context,
									   CL_MEM_READ_WRITE,
									   4 * sizeof(float) * numParticles,
									   NULL,
									   &error);
	clb_check_error(error, "clCreateBuffer");
	
		
	particleProgram = [[ComputeProgram alloc] initWithContext: simpleCompute.context
													   deviceID: simpleCompute.deviceID
														   file: [[NSBundle mainBundle] pathForResource:@"particles" ofType:@"cl"]
													 kernelName: @"particleStep"
												   buildOptions:@""];
	
	particleShader = [[ShaderProgram alloc] initWithVertexShader:[[NSBundle mainBundle] pathForResource:@"particle" ofType:@"vsh"]
												  fragmentShader:[[NSBundle mainBundle] pathForResource:@"particle" ofType:@"fsh"]
												  geometryShader:[[NSBundle mainBundle] pathForResource:@"particle" ofType:@"gsh"]];
	
	
	billboard = [[GLTexture alloc] initWithFile:
				 [[NSBundle mainBundle] pathForResource:@"particle" ofType:@"png"]
										  flags: kGLTextureGenerateMipMaps]; 
		
	[particleShader setInputType: GL_POINTS];
	[particleShader setOutputType: GL_TRIANGLE_STRIP];
	[particleShader setNumVerticesOut: 4];	
	
	[particleShader linkProgram];
	rb_check_error_simple();

	
	rb_check_error_simple();
	
	[self stepParticleSystem: 0.0f reset: 1 modelviewMatrix: rb_mat4_identity()];
	
	int i;
	for (i=0; i<4000; i++) {
		[self stepParticleSystem: 0.01f reset: 0 modelviewMatrix: rb_mat4_identity()];
	}
	
	return self;
	
}

-(void)stepParticleSystem:(float)timeStep reset:(BOOL)reset modelviewMatrix:(mat4)modelview {
	
	
	/*
	 __global float *positions,
	 __global float *velocities,
	 const float timeStep,
	 const float animationTime,
	 const float lifeTime,
	 const int resetSignal,
	 const float origin
	 */
			
	cl_int error = CL_SUCCESS;
	
	float origin[4] = { -0.15f, -0.2f, -0.4f, 1.0f };
	
	float lifeTime = 3.0f;
	
	int resetSignal = reset;
	
	if ( resetSignal ) {
		animationTime = 0.0f;
	}
	else {
		animationTime += timeStep;
	}
			
	cl_command_queue queue = simpleCompute.queue;
	
	error |= clEnqueueAcquireGLObjects(queue, 1, &clPositionsBuffer, 0, 0, 0);
	error |= clEnqueueAcquireGLObjects(queue, 1, &clVelocitiesBuffer, 0, 0, 0);
	error |= clEnqueueAcquireGLObjects(queue, 1, &clAgeBuffer, 0, 0, 0);

	clb_check_error_simple(error);
	
	cl_kernel kernel = particleProgram.kernel;
	clb_check_error_simple(clSetKernelArg(kernel, 0,  sizeof(cl_mem), &clPositionsBuffer));
	clb_check_error_simple(clSetKernelArg(kernel, 1,  sizeof(cl_mem), &clVelocitiesBuffer));
	clb_check_error_simple(clSetKernelArg(kernel, 2,  sizeof(cl_mem), &clAgeBuffer));
	clb_check_error_simple(clSetKernelArg(kernel, 3,  sizeof(float), &timeStep));
	clb_check_error_simple(clSetKernelArg(kernel, 4,  sizeof(float), &animationTime));
	clb_check_error_simple(clSetKernelArg(kernel, 5,  sizeof(float), &lifeTime));
	clb_check_error_simple(clSetKernelArg(kernel, 6,  sizeof(int), &resetSignal));
	clb_check_error_simple(clSetKernelArg(kernel, 7,  4.0f * sizeof(float), &origin));
	clb_check_error_simple(clSetKernelArg(kernel, 8,  sizeof(int), &numParticles));

	clb_check_error_simple(clSetKernelArg(kernel, 9,  sizeof(cl_mem), &eyespaceParticles));
	
	float mv1[4], mv2[4], mv3[4], mv4[4];
	mv1[0] = modelview.entries[0];
	mv1[1] = modelview.entries[1];
	mv1[2] = modelview.entries[2];
	mv1[3] = modelview.entries[3];
	
	mv2[0] = modelview.entries[4];
	mv2[1] = modelview.entries[5];
	mv2[2] = modelview.entries[6];
	mv2[3] = modelview.entries[7];
	
	mv3[0] = modelview.entries[8];
	mv3[1] = modelview.entries[9];
	mv3[2] = modelview.entries[10];
	mv3[3] = modelview.entries[11];
	
	mv4[0] = modelview.entries[12];
	mv4[1] = modelview.entries[13];
	mv4[2] = modelview.entries[14];
	mv4[3] = modelview.entries[15];
	
	clb_check_error_simple(clSetKernelArg(kernel, 10, 4.0f * sizeof(float), &mv1 ));
	clb_check_error_simple(clSetKernelArg(kernel, 11, 4.0f * sizeof(float), &mv2 ));
	clb_check_error_simple(clSetKernelArg(kernel, 12, 4.0f * sizeof(float), &mv3 ));
	clb_check_error_simple(clSetKernelArg(kernel, 13, 4.0f * sizeof(float), &mv4 ));

	int work_dim = 1;
	size_t local_work_size = 64;
	size_t global_work_size = ((numParticles + local_work_size - 1) / numParticles) * numParticles;
		
	cl_event event;
	error = clEnqueueNDRangeKernel(simpleCompute.queue, kernel, work_dim, NULL, &global_work_size, &local_work_size, 0, NULL, &event);
	clb_check_error_simple(error);
			
	error |= clEnqueueReleaseGLObjects(queue, 1, &clPositionsBuffer, 0, 0, 0);
	error |= clEnqueueReleaseGLObjects(queue, 1, &clVelocitiesBuffer, 0, 0, 0);
	error |= clEnqueueReleaseGLObjects(queue, 1, &clAgeBuffer, 0, 0, 0);
	
	//clFinish(queue);
	//clb_print_event_profiling_info(event, "particle system");
	
	clb_check_error_simple(error);	
			
}

-(void)drawParticlesWithModelviewMatrix:(mat4)modelview projectionMatrix:(mat4)projection {
		
	[particleShader useProgram];
	rb_check_error_simple();

	glPointSize(4.5f);
	
	glEnable(GL_BLEND);
	glBlendFunc(GL_SRC_ALPHA, GL_ONE);
	glDepthFunc(GL_LEQUAL);
	
	glUniformMatrix4fv( [particleShader locationForUniform:@"projectionMatrix"],
					   1, 0, projection.entries );
	glUniformMatrix4fv( [particleShader locationForUniform:@"modelViewMatrix"],
					   1, 0, modelview.entries );

	glActiveTexture(GL_TEXTURE0);
	glBindTexture(GL_TEXTURE_2D, billboard.descriptor);
	
	rb_check_error_simple();

	GLuint samplerIndex = [particleShader getLocationForAttribute:@"sampler"];
	glUniform1i(samplerIndex, 0);
	
	GLuint ageIndex = [particleShader getLocationForAttribute:@"age_in"];
	rb_check_error_simple();
	glBindBuffer(GL_ARRAY_BUFFER, glAgeBuffer);
	glEnableVertexAttribArray(ageIndex);
	glVertexAttribPointer(ageIndex, 1, GL_FLOAT, 0, 0, 0);		
	
	GLuint positionIndex = [particleShader getLocationForAttribute:@"position"];
	rb_check_error_simple();

	glBindBuffer(GL_ARRAY_BUFFER, glPositionsBuffer);
	glEnableVertexAttribArray(positionIndex);
	glVertexAttribPointer(positionIndex, 4, GL_FLOAT, 0, 0, 0);	
	rb_check_error_simple();

	glDrawArrays(GL_POINTS, 0, numParticles);
	
	glBindBuffer(GL_ARRAY_BUFFER, 0);
	glDisableVertexAttribArray(positionIndex);
	glDisableVertexAttribArray(ageIndex);

	rb_check_error_simple();
	
	glActiveTexture(GL_TEXTURE0);
	glBindTexture(GL_TEXTURE_2D, 0);
	
	glEnable(GL_DEPTH_TEST);
	glDisable(GL_BLEND);
}

@end
